#include "CudaEpiModel.h"
#include "CudaParam.h"
#include "CudaPerson.h"
#include <thrust/copy.h>
#include <thrust/binary_search.h>
#include <thrust/count.h>
#include <thrust/sort.h>
#include <thrust/remove.h>
#include <sys/time.h> 
#include <thrust/sequence.h>
#include <thrust/partition.h>
#include <omp.h>

dCPUtoGPUPersonVec* cpu_to_gpu_device_temp;
hCPUtoGPUPersonVec* cpu_to_gpu_host_temp;
dGPUtoCPUPersonVec* gpu_to_cpu_device_temp;
hGPUtoCPUPersonVec* gpu_to_cpu_host_temp;


thrust::host_vector<dCommunity> h_comm_vec;

//thrust::device_vector<bool>* infect_result;


bool compareModel(EpiModel& compare_model){
	return true;
}

__global__ void tester(dCommunity* comm_test, int size){
	for(int i = 0; i < size; i++){
		dCommunity& temp = comm_test[i];
		for(int j = 0; j < TAG; j++){
			printf("community %d age %d has %d\n",i,j,temp.ninf[j]);
		}
	}
}

void cudaDay(){
	// copy person to deivce
	struct timeval start, finish;
	double t1,t2;
	
	 //gettimeofday(&start,NULL);

	/*cudaMemcpy(thrust::raw_pointer_cast(ppl_vec->data()),&model->pvec[0],
		(model->nNumPerson)*sizeof(Person),cudaMemcpyHostToDevice);*/
	
	CopyFromCPUtoGPU(model->pvec,ppl_vec,cpu_to_gpu_host_temp,cpu_to_gpu_device_temp,model->nNumPerson);
	
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "transfer 11 taken is: " << (t2 - t1) << " secs" << endl;*/
	
	//gettimeofday(&start,NULL);
	cudaMemcpy(thrust::raw_pointer_cast(tract_vec->data()),&model->tractvec[0],
		(model->nNumPerson)*sizeof(Tract),cudaMemcpyHostToDevice);
	//err = cudaThreadSynchronize();
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "transfer 12 taken is: " << (t2 - t1) << " secs" << endl;*/
	
	//gettimeofday(&start,NULL);
	copyCommunity(true);
	//err = cudaThreadSynchronize();
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "transfer 13 taken is: " << (t2 - t1) << " secs" << endl;*/
	
	
	// #pragma omp parallel for
	// for(int i = 0; i < 3; i++){
		// switch(i){
		// case 0:
			// cudaMemcpy(
				// thrust::raw_pointer_cast(ppl_vec->data()),&model->pvec[0],
				// (model->nNumPerson)*sizeof(Person),cudaMemcpyHostToDevice);
			// break;
		// case 1:
			// cudaMemcpy(
				// thrust::raw_pointer_cast(ppl_vec->data()),&model->pvec[0],
				// (model->nNumPerson)*sizeof(Person),cudaMemcpyHostToDevice);
			// break;
		// case 2:
			// copyCommunity(true);
			// break;
		// }
	// }
	
	
	//gettimeofday(&start,NULL);
	thrust::for_each(ppl_vec->begin(),ppl_vec->end(),
		DayOp1(model->nTimer,model->AVEs,model->AVEi));
		
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "comp 1 taken is: " << (t2 - t1) << " secs" << endl;*/
	
	// thrust::for_each(comm_visitor_vec->begin(),comm_visitor_vec->end(),
		// DayOp2(model->AVEs,model->AVEi));

	// err = cudaThreadSynchronize();
        // if(err != cudaSuccess) cout << "cudaThreadSynchronize error: " << cudaGetErrorString(err) << endl;
		
	// for(int i = 0; i < model->nNumCommunities; i++){
		// Community& temp = model->commvec[i];
		// cout << "Community " << i << " has ppl: " << temp.nLastPerson - temp.nFirstPerson << 
		// " and workers: " << temp.workers.size() << endl;
	// }

	// thrust::for_each(comm_vec->begin(),comm_vec->end(),
		// DayOp3(thrust::raw_pointer_cast(ppl_vec->data()),
				// thrust::raw_pointer_cast(tract_vec->data()),
				// thrust::raw_pointer_cast(comm_vec->data()),
				// thrust::raw_pointer_cast(rand_vec->data()),
				// model->bTravel,
				// thrust::raw_pointer_cast(comm_worker_vec->data()),
				// thrust::raw_pointer_cast(comm_worker_start_vec->data()),
				// thrust::raw_pointer_cast(comm_worker_end_vec->data()),
				// thrust::raw_pointer_cast(comm_visitor_vec->data()),
				// thrust::raw_pointer_cast(comm_visitor_start_vec->data()),
				// thrust::raw_pointer_cast(comm_visitor_end_vec->data()),
				// thrust::raw_pointer_cast(infect_result.data()))
				// );
	
	//gettimeofday(&start,NULL);
	thrust::device_vector<int> infect_opt(model->nNumPerson);
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "alocation time is: " << (t2 - t1) << " secs" << endl;*/

	//gettimeofday(&start,NULL);
	thrust::sequence(infect_opt.begin(),infect_opt.end());
	/*err = cudaThreadSynchronize();
	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "sequence time is: " << (t2 - t1) << " secs" << endl;*/
	
	//gettimeofday(&start,NULL);
	thrust::device_vector<int>::iterator infect_opt_end2 = 
		thrust::partition(infect_opt.begin(),infect_opt.end(),
			IsInfectiousDay(thrust::raw_pointer_cast(ppl_vec->data())));
	//err = cudaThreadSynchronize();
	//gettimeofday(&finish,NULL);
	//t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	//t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	//cout << "partition time 1 is: " << (t2 - t1) << " secs" << endl;
	//cout << "partition size 1 is: " << infect_opt_end2 - infect_opt.begin() << endl;
		
	//gettimeofday(&start,NULL);
	thrust::device_vector<int>::iterator infect_opt_end1 = 
		thrust::partition(infect_opt.begin(),infect_opt_end2,
			IsHomeWorkerDay(thrust::raw_pointer_cast(ppl_vec->data())));

	/*err = cudaThreadSynchronize();
	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "partition time 2 is: " << (t2 - t1) << " secs" << endl;
	cout << "partition size 2 is: " << infect_opt_end1 - infect_opt.begin() << endl;*/

	// gettimeofday(&finish,NULL);
	/*t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "partition 1 taken is: " << (t2 - t1) << " secs" << endl;*/
	
	// cout << "Num of infecting people: " << (infect_opt_end - infect_opt->begin()) << endl;
	
	thrust::device_vector<bool> infect_result(model->nNumPerson,false);

	cout << "first group: " << infect_opt_end1 - infect_opt.begin() << endl;
	gettimeofday(&start,NULL);

	int numBlks = infect_opt_end1 - infect_opt.begin();
	
	if(numBlks != 0){
		int numThreads = model->nNumPerson / numBlks;

		//if(numThreads > 1024) numThreads = 1024;
		
		if(numThreads > 256)  numThreads = 256;
		else if(numThreads > 128)  numThreads = 128;
		else if(numThreads > 64) numThreads = 64;
		else if(numThreads > 32) numThreads = 32;
		else if(numThreads > 16) numThreads = 16;
		else if(numThreads > 8) numThreads = 8;
		
		DayOp3EvenBetter<<<numBlks,numThreads>>>(
				thrust::raw_pointer_cast(infect_opt.data()),
				thrust::raw_pointer_cast(ppl_vec->data()),
				thrust::raw_pointer_cast(tract_vec->data()),
				thrust::raw_pointer_cast(comm_vec->data()),
				thrust::raw_pointer_cast(rand_vec->data()),
				model->bTravel,
				thrust::raw_pointer_cast(comm_worker_vec->data()),
				thrust::raw_pointer_cast(comm_worker_start_vec->data()),
				thrust::raw_pointer_cast(comm_worker_end_vec->data()),
				thrust::raw_pointer_cast(comm_visitor_vec->data()),
				thrust::raw_pointer_cast(comm_visitor_start_vec->data()),
				thrust::raw_pointer_cast(comm_visitor_end_vec->data()),
				thrust::raw_pointer_cast(infect_result.data()),true);
	}
	err = cudaThreadSynchronize();
	if(err != cudaSuccess)
	{
		// print the CUDA error message and exit
		printf("CUDA error1: %s\n", cudaGetErrorString(err));
		getchar();
	}

	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "computing 1 is: " << (t2 - t1) << " secs" << endl;

	
	cout << "second group: " << infect_opt_end2 - infect_opt_end1 << endl;
	 gettimeofday(&start,NULL);

	 numBlks = infect_opt_end2 - infect_opt_end1;
	
	 thrust::sort(infect_opt_end1,infect_opt_end2,
		SortPplByDayComm(thrust::raw_pointer_cast(ppl_vec->data())));

	 if(numBlks != 0){
		int numThreads = model->nNumPerson / numBlks;

		//if(numThreads > 1024) numThreads = 1024;

		
		if(numThreads > 256)  numThreads = 256;
		else if(numThreads > 128)  numThreads = 128;
		else if(numThreads > 64) numThreads = 64;
		else if(numThreads > 32) numThreads = 32;
		else if(numThreads > 16) numThreads = 16;
		else if(numThreads > 8) numThreads = 8;
		
		DayOp3EvenBetter<<<numBlks,numThreads>>>(
				thrust::raw_pointer_cast(infect_opt.data()) + (infect_opt_end1 - infect_opt.begin()),
				thrust::raw_pointer_cast(ppl_vec->data()),
				thrust::raw_pointer_cast(tract_vec->data()),
				thrust::raw_pointer_cast(comm_vec->data()),
				thrust::raw_pointer_cast(rand_vec->data()),
				model->bTravel,
				thrust::raw_pointer_cast(comm_worker_vec->data()),
				thrust::raw_pointer_cast(comm_worker_start_vec->data()),
				thrust::raw_pointer_cast(comm_worker_end_vec->data()),
				thrust::raw_pointer_cast(comm_visitor_vec->data()),
				thrust::raw_pointer_cast(comm_visitor_start_vec->data()),
				thrust::raw_pointer_cast(comm_visitor_end_vec->data()),
				thrust::raw_pointer_cast(infect_result.data()),false);
	 }


	 err = cudaThreadSynchronize();
	 if(err != cudaSuccess)
	 {
		 // print the CUDA error message and exit
		 printf("CUDA error2: %s\n", cudaGetErrorString(err));
		 getchar();
	 }
	
	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "computing 2 is: " << (t2 - t1) << " secs" << endl;
	

	//gettimeofday(&start,NULL);
	thrust::sequence(infect_opt.begin(),infect_opt.end());
	//gettimeofday(&finish,NULL);
	//t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	//t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	//cout << "sequence 2 taken is: " << (t2 - t1) << " secs" << endl;

	//gettimeofday(&start,NULL);
	thrust::device_vector<int>::iterator infected_end_itr = 
		thrust::partition(infect_opt.begin(),infect_opt.end(),
			IsInfectedDay(thrust::raw_pointer_cast(infect_result.data())));

	/*err = cudaThreadSynchronize();
	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "partition 3 taken is: " << (t2 - t1) << " secs" << endl;*/
	
	//gettimeofday(&start,NULL);
	thrust::for_each(infect_opt.begin(),infected_end_itr,
		InfectOp(
				thrust::raw_pointer_cast(rand_vec->data()),
				model->nTimer,model->AVEp,
				model->fSymptomaticAscertainment,model->bTrigger,
				model->fLiberalLeaveCompliance,model->fIsolationCompliance,
				thrust::raw_pointer_cast(comm_vec->data()),
				thrust::raw_pointer_cast(ppl_vec->data())
				)
				);
	/*err = cudaThreadSynchronize();
	gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "infection taken is: " << (t2 - t1) << " secs" << endl;*/

	
	//gettimeofday(&start,NULL);
	/*cudaMemcpy(
		&model->pvec[0],thrust::raw_pointer_cast(ppl_vec->data()),
		(model->nNumPerson)*sizeof(Person),cudaMemcpyDeviceToHost);*/
	CopyFromGPUtoCPU(model->pvec,ppl_vec,gpu_to_cpu_host_temp,gpu_to_cpu_device_temp,model->nNumPerson);
	
	
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "transfer 21 taken is: " << (t2 - t1) << " secs" << endl;	*/
	
	/*gettimeofday(&start,NULL);*/
	copyCommunity(false);
	/*gettimeofday(&finish,NULL);
	t1 = (double)start.tv_sec + ((double)start.tv_usec/1000000.0);
	t2 = (double)finish.tv_sec + ((double)finish.tv_usec/1000000.0);
	cout << "transfer 22 taken is: " << (t2 - t1) << " secs" << endl;	*/
}


// true: transfer community vector from host to device
// false: transfer community vector from device to host
void copyCommunity(bool host_to_device){
	if(host_to_device){
		// #pragma omp parallel for
		for(int i = 0; i < model->nNumCommunities; i++){
			h_comm_vec[i].copyFromCommunity(model->commvec[i]);
		}
		cudaMemcpy(thrust::raw_pointer_cast(comm_vec->data()),
			thrust::raw_pointer_cast(h_comm_vec.data()),model->nNumCommunities*sizeof(dCommunity),
			cudaMemcpyHostToDevice);
	
	}else{
		cudaMemcpy(thrust::raw_pointer_cast(h_comm_vec.data()),
			thrust::raw_pointer_cast(comm_vec->data()),model->nNumCommunities*sizeof(dCommunity),
			cudaMemcpyDeviceToHost);
		// #pragma omp parallel for
		for(int i = 0; i < model->nNumCommunities; i++){
			h_comm_vec[i].copyToCommunity(model->commvec[i]); 
		}
		
	}
}

void clean(){
	delete ppl_vec;
	delete comm_vec;
	delete comm_worker_vec; 
	delete comm_worker_start_vec; 
	delete comm_worker_end_vec;  

	delete comm_visitor_vec; 
	delete comm_visitor_start_vec; 
	delete comm_visitor_end_vec; 

	delete comm_worker_counter;
	delete comm_visitor_counter;
	delete rand_vec;
}


void cudaInitialize(EpiModel& input_model){
	model = &input_model;
    // transfer constant  information to constant memory
	cudaMemcpyToSymbol(dvload,model->vload,VLOADNSUB*VLOADNDAY*sizeof(double)); // vload
    cudaMemcpyToSymbol(dseasonality,model->seasonality,MAXRUNLENGTH*sizeof(double)); // seasonality
    cudaMemcpyToSymbol(dfVaccineEfficacyByAge,model->fVaccineEfficacyByAge,TAG*sizeof(double)); // fVaccineEfficacyByAge
    cudaMemcpyToSymbol(dVaccineData,model->VaccineData,NUMVACCINES*sizeof(vaccinedatastruct)); // vaccineData
	cudaMemcpyToSymbol(dcpfc,cpfc,TAG*sizeof(double));
	cudaMemcpyToSymbol(dcpfa,cpfa,TAG*sizeof(double));
	cudaMemcpyToSymbol(dcphcc,cphcc,TAG*sizeof(double));
	cudaMemcpyToSymbol(dcphca,cphca,TAG*sizeof(double));
	cudaMemcpyToSymbol(dwithdrawcdf32,model->withdrawcdf32,3*WITHDRAWDAYS*sizeof(unsigned int)); // withdrawcef32
	cudaMemcpyToSymbol(dincubationcdf32,incubationcdf32,3*sizeof(unsigned int)); // incubationcdf32
	cudaMemcpyToSymbol(dcps,cps,10*sizeof(double)); // cps
	cudaMemcpyToSymbol(d_travel_length_cdf,travel_length_cdf,12*sizeof(double)); // travel_length_cdf
	cudaMemcpyToSymbol(d_travel_pr,travel_pr,TAG*sizeof(double)); // travel_pr
	cudaMemcpyToSymbol(dnSchoolOpeningDays,model->nSchoolOpeningDays,56*sizeof(int));
	cudaMemcpyToSymbol(dnVaccinePriorities2,model->nVaccinePriorities2,PRIORITY_LAST*sizeof(unsigned char));
	
	ppl_vec = new dPersonVec(model->nNumPerson);
	cudaMemcpy(thrust::raw_pointer_cast(ppl_vec->data()),&model->pvec[0],
		(model->nNumPerson)*sizeof(Person),cudaMemcpyHostToDevice);
	err = cudaThreadSynchronize();
	
	cpu_to_gpu_device_temp = new dCPUtoGPUPersonVec(model->nNumPerson);
	cpu_to_gpu_host_temp = new hCPUtoGPUPersonVec(model->nNumPerson);
	gpu_to_cpu_device_temp = new dGPUtoCPUPersonVec(model->nNumPerson);
	gpu_to_cpu_host_temp = new hGPUtoCPUPersonVec(model->nNumPerson);
	
	// copy community vector
	comm_vec = new dCommunityVec(model->nNumCommunities);
	h_comm_vec.reserve(model->nNumCommunities);

	// copy tract vector
	tract_vec = new dTractVec(model->nNumTractsTotal);

	rand_vec = new dRandVec(model->nNumPerson);
	srand(100);
	hRandVec rand_temp;
	for(int i = 0; i < model->nNumPerson; i++){rand_temp.push_back(RandGen(rand(),rand(),rand(),rand()));}
	*rand_vec = rand_temp;

	// initialize community counter and construct comm_woker_vec and comm_visitior_vec
	unsigned int worker_size = 0; 
	unsigned int visitor_size = 0;
	hCommWorkerVec comm_worker_temp,comm_worker_counter_temp;
	hCommVisitorVec comm_visitor_temp,comm_visitor_counter_temp;

	Person dummy_ppl;
	for(unsigned int i = 0; i < model->nNumCommunities; i++){
		comm_worker_counter_temp.push_back(dCommWorker(i,0));
		comm_visitor_counter_temp.push_back(dCommVisitor(i,dummy_ppl));

		for(unsigned int j = 0; j < model->commvec[i].workers.size(); j++){
			comm_worker_temp.push_back(dCommWorker(i,model->commvec[i].workers[j]));
		}
		worker_size += model->commvec[i].workers.size();	

		for(list< Person >::iterator it = model->commvec[i].visitors.begin();
		    it != model->commvec[i].visitors.end(); ++it){
			Person &p = *it;
			comm_visitor_temp.push_back(dCommVisitor(i,p));
			visitor_size++;
		}	
	}

	comm_worker_counter = new dCommWorkerVec(model->nNumCommunities);
	comm_visitor_counter = new dCommVisitorVec(model->nNumCommunities);
	*comm_worker_counter = comm_worker_counter_temp;
	*comm_visitor_counter = comm_visitor_counter_temp;

	comm_worker_vec = new dCommWorkerVec(worker_size);
	comm_visitor_vec = new dCommVisitorVec(visitor_size);
	*comm_worker_vec = comm_worker_temp;
	*comm_visitor_vec = comm_visitor_temp;

	// find the start index and end index of worker and visitor for each community
	comm_worker_start_vec = new dIdxVec(model->nNumCommunities);
	comm_worker_end_vec = new dIdxVec(model->nNumCommunities);
	// since visitor is initially empty, so no need to setup
	comm_visitor_start_vec = new dIdxVec(model->nNumCommunities);
	comm_visitor_end_vec = new dIdxVec(model->nNumCommunities);

	
	thrust::lower_bound(
		comm_worker_vec->begin(),comm_worker_vec->end(),
		comm_worker_counter->begin(),comm_worker_counter->end(),
		comm_worker_start_vec->begin());

	thrust::upper_bound(
		comm_worker_vec->begin(),comm_worker_vec->end(),
		comm_worker_counter->begin(),comm_worker_counter->end(),
		comm_worker_end_vec->begin());


	cudaMalloc((void**)&nNumTAPDone,sizeof(unsigned int));
	cudaMalloc((void**)&nNumWantAV,sizeof(unsigned int));
	cudaMalloc((void**)&nNumWantVaccine,sizeof(unsigned int));
	cudaMalloc((void**)&nNumAntiviralsUsed,sizeof(unsigned int));
	
	cudaMemcpy(nNumTAPDone,&model->nNumTAPDone,sizeof(unsigned int),cudaMemcpyHostToDevice);
	cudaMemcpy(nNumWantAV,&model->nNumWantAV,sizeof(unsigned int),cudaMemcpyHostToDevice);
	cudaMemcpy(nNumWantVaccine,&model->nNumWantVaccine,sizeof(unsigned int),cudaMemcpyHostToDevice);
	cudaMemcpy(nNumAntiviralsUsed,&model->nNumAntiviralsUsed,sizeof(unsigned int),cudaMemcpyHostToDevice);
	
	//infect_result = new thrust::device_vector<bool>(model->nNumPerson,false);
}
